//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-31968024
// Cuda compilation tools, release 12.0, V12.0.76
// Based on NVVM 7.0.1
//

.version 8.0
.target sm_52
.address_size 64

	// .globl	__closesthit__oxMain
.const .align 16 .b8 params[384];

.visible .entry __closesthit__oxMain()
{
	.reg .pred 	%p<30>;
	.reg .b16 	%rs<12>;
	.reg .f32 	%f<245>;
	.reg .b32 	%r<66>;
	.reg .b64 	%rd<22>;


	// begin inline asm
	call (%r3), _optix_read_primitive_idx, ();
	// end inline asm
	// begin inline asm
	call (%f58, %f59), _optix_get_triangle_barycentrics, ();
	// end inline asm
	mov.u32 	%r15, 0;
	// begin inline asm
	call (%rd3), _optix_call_direct_callable,(%r15);
	// end inline asm
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b32 param0;
	st.param.b32 	[param0+0], %r3;
	.param .align 8 .b8 param1[8];
	st.param.v2.f32 	[param1+0], {%f58, %f59};
	.param .align 8 .b8 retval0[32];
	prototype_0 : .callprototype (.param .align 8 .b8 _[32]) _ (.param .b32 _, .param .align 8 .b8 _[8]);
	call (retval0), 
	%rd3, 
	(
	param0, 
	param1
	)
	, prototype_0;
	ld.param.v2.f32 	{%f1, %f2}, [retval0+0];
	ld.param.f32 	%f3, [retval0+8];
	ld.param.v4.b8 	{%rs1, %rs2, %rs3, %rs4}, [retval0+12];
	ld.param.v2.f32 	{%f4, %f5}, [retval0+16];
	ld.param.f32 	%f6, [retval0+24];
	ld.param.v4.b8 	{%rs5, %rs6, %rs7, %rs8}, [retval0+28];
	} // callseq 0
	mov.u32 	%r17, 1;
	// begin inline asm
	call (%f60), _optix_get_ray_tmax, ();
	// end inline asm
	mov.u32 	%r13, 3;
	// begin inline asm
	call (%r5), _optix_get_payload, (%r13);
	// end inline asm
	mov.b32 	%f8, %r5;
	mov.u32 	%r8, 4;
	// begin inline asm
	call (%r7), _optix_get_payload, (%r8);
	// end inline asm
	mov.b32 	%f9, %r7;
	mov.u32 	%r10, 5;
	// begin inline asm
	call (%r9), _optix_get_payload, (%r10);
	// end inline asm
	mov.b32 	%f10, %r9;
	mov.u32 	%r12, 6;
	// begin inline asm
	call (%r11), _optix_get_payload, (%r12);
	// end inline asm
	mov.b32 	%f11, %r11;
	mov.b32 	%r14, %f60;
	// begin inline asm
	call _optix_set_payload, (%r13, %r14);
	// end inline asm
	mov.b32 	%r16, %f1;
	// begin inline asm
	call _optix_set_payload, (%r15, %r16);
	// end inline asm
	mov.b32 	%r18, %f2;
	// begin inline asm
	call _optix_set_payload, (%r17, %r18);
	// end inline asm
	mov.b32 	%r20, %f3;
	mov.u32 	%r19, 2;
	// begin inline asm
	call _optix_set_payload, (%r19, %r20);
	// end inline asm
	// begin inline asm
	call (%f61), _optix_get_world_ray_direction_x, ();
	// end inline asm
	// begin inline asm
	call (%f62), _optix_get_world_ray_direction_y, ();
	// end inline asm
	// begin inline asm
	call (%f63), _optix_get_world_ray_direction_z, ();
	// end inline asm
	mul.f32 	%f64, %f2, %f62;
	fma.rn.f32 	%f65, %f1, %f61, %f64;
	fma.rn.f32 	%f12, %f3, %f63, %f65;
	setp.ge.f32 	%p2, %f12, 0f00000000;
	ld.const.f32 	%f66, [params+372];
	setp.eq.f32 	%p3, %f66, 0f00000000;
	and.pred  	%p4, %p3, %p2;
	@%p4 bra 	$L__BB0_23;
	bra.uni 	$L__BB0_1;

$L__BB0_23:
	// begin inline asm
	call _optix_set_payload, (%r8, %r15);
	// end inline asm
	// begin inline asm
	call _optix_set_payload, (%r10, %r15);
	// end inline asm
	// begin inline asm
	call _optix_set_payload, (%r12, %r15);
	// end inline asm
	bra.uni 	$L__BB0_24;

$L__BB0_1:
	div.rn.f32 	%f67, %f4, 0f41200000;
	cvt.rzi.s32.f32 	%r1, %f67;
	setp.lt.s32 	%p5, %r1, 0;
	@%p5 bra 	$L__BB0_5;

	cvt.rzi.s32.f32 	%r21, %f4;
	cvt.rn.f32.s32 	%f68, %r21;
	sub.f32 	%f15, %f4, %f68;
	ld.const.u64 	%rd4, [params+248];
	cvta.to.global.u64 	%rd5, %rd4;
	mul.wide.s32 	%rd6, %r1, 8;
	add.s64 	%rd7, %rd5, %rd6;
	ld.global.s32 	%rd8, [%rd7];
	tex.2d.v4.f32.f32 	{%f69, %f70, %f71, %f72}, [%rd8, {%f15, %f5}];
	add.f32 	%f73, %f69, 0fB4000000;
	mov.f32 	%f74, 0f00000000;
	max.f32 	%f234, %f73, %f74;
	add.f32 	%f75, %f70, 0fB4000000;
	max.f32 	%f235, %f75, %f74;
	add.f32 	%f76, %f71, 0fB4000000;
	max.f32 	%f236, %f76, %f74;
	ld.const.u64 	%rd9, [params+256];
	cvta.to.global.u64 	%rd10, %rd9;
	add.s64 	%rd11, %rd10, %rd6;
	ld.global.u64 	%rd1, [%rd11];
	cvt.u32.u64 	%r22, %rd1;
	setp.lt.s32 	%p6, %r22, 1;
	@%p6 bra 	$L__BB0_4;

	cvt.s64.s32 	%rd12, %rd1;
	tex.2d.v4.u32.f32 	{%r23, %r24, %r25, %r26}, [%rd12, {%f15, %f5}];
	cvt.u16.u32 	%rs9, %r23;
	cvt.u16.u32 	%rs10, %r24;
	cvt.u16.u32 	%rs11, %r25;
	// begin inline asm
	{  cvt.f32.f16 %f77, %rs9;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f78, %rs10;}

	// end inline asm
	// begin inline asm
	{  cvt.f32.f16 %f79, %rs11;}

	// end inline asm
	ld.const.f32 	%f80, [params+328];
	mul.f32 	%f81, %f80, %f77;
	mul.f32 	%f82, %f80, %f78;
	mul.f32 	%f83, %f80, %f79;
	mul.f32 	%f84, %f81, 0f437F0000;
	cvt.rmi.f32.f32 	%f85, %f84;
	mul.f32 	%f86, %f82, 0f437F0000;
	cvt.rmi.f32.f32 	%f87, %f86;
	mul.f32 	%f88, %f83, 0f437F0000;
	cvt.rmi.f32.f32 	%f89, %f88;
	add.f32 	%f234, %f234, %f85;
	add.f32 	%f235, %f235, %f87;
	add.f32 	%f236, %f236, %f89;

$L__BB0_4:
	selp.f32 	%f90, 0fBF800000, 0f3F800000, %p2;
	mul.f32 	%f91, %f90, %f234;
	mul.f32 	%f92, %f91, %f9;
	setp.gt.f32 	%p8, %f60, %f8;
	selp.f32 	%f93, %f92, %f91, %p8;
	mul.f32 	%f94, %f235, %f10;
	selp.f32 	%f95, %f94, %f235, %p8;
	mul.f32 	%f96, %f236, %f11;
	selp.f32 	%f97, %f96, %f236, %p8;
	mov.b32 	%r28, %f93;
	mov.u32 	%r27, 4;
	// begin inline asm
	call _optix_set_payload, (%r27, %r28);
	// end inline asm
	mov.b32 	%r30, %f95;
	mov.u32 	%r29, 5;
	// begin inline asm
	call _optix_set_payload, (%r29, %r30);
	// end inline asm
	mov.b32 	%r32, %f97;
	mov.u32 	%r31, 6;
	// begin inline asm
	call _optix_set_payload, (%r31, %r32);
	// end inline asm
	bra.uni 	$L__BB0_24;

$L__BB0_5:
	ld.const.u64 	%rd13, [params+224];
	cvta.to.global.u64 	%rd14, %rd13;
	neg.s32 	%r33, %r1;
	cvt.s64.s32 	%rd2, %r33;
	mul.wide.s32 	%rd15, %r33, 16;
	add.s64 	%rd16, %rd14, %rd15;
	ld.global.v4.f32 	{%f98, %f99, %f100, %f240}, [%rd16];
	mov.b32 	%r34, %f98;
	shr.s32 	%r35, %r34, 16;
	cvt.rn.f32.s32 	%f102, %r35;
	div.rn.f32 	%f103, %f102, 0f437F0000;
	mul.f32 	%f237, %f99, %f103;
	bfe.u32 	%r36, %r34, 8, 8;
	cvt.rn.f32.s32 	%f104, %r36;
	div.rn.f32 	%f105, %f104, 0f437F0000;
	mul.f32 	%f238, %f99, %f105;
	and.b32  	%r37, %r34, 255;
	cvt.rn.f32.s32 	%f106, %r37;
	div.rn.f32 	%f107, %f106, 0f437F0000;
	mul.f32 	%f239, %f99, %f107;
	setp.gtu.f32 	%p9, %f240, 0f00000000;
	@%p9 bra 	$L__BB0_7;

	neg.f32 	%f240, %f240;
	ld.const.u64 	%rd17, [params+264];
	cvta.to.global.u64 	%rd18, %rd17;
	shl.b64 	%rd19, %rd2, 3;
	add.s64 	%rd20, %rd18, %rd19;
	mul.lo.s32 	%r38, %r1, -10;
	cvt.rn.f32.s32 	%f108, %r38;
	neg.f32 	%f109, %f4;
	sub.f32 	%f110, %f109, %f108;
	ld.global.s32 	%rd21, [%rd20];
	mul.f32 	%f111, %f110, 0f41200000;
	mul.f32 	%f112, %f5, 0f41200000;
	tex.2d.v4.f32.f32 	{%f113, %f114, %f115, %f116}, [%rd21, {%f111, %f112}];
	mul.f32 	%f237, %f237, %f113;
	mul.f32 	%f238, %f238, %f114;
	mul.f32 	%f239, %f239, %f115;

$L__BB0_7:
	ld.const.u32 	%r2, [params+340];
	setp.lt.s32 	%p10, %r2, 1;
	mov.f32 	%f244, 0f3F800000;
	@%p10 bra 	$L__BB0_22;

	mul.f32 	%f42, %f60, %f240;
	mov.f32 	%f122, 0f40800000;
	abs.f32 	%f44, %f42;
	setp.lt.f32 	%p11, %f44, 0f00800000;
	mul.f32 	%f124, %f44, 0f4B800000;
	selp.f32 	%f125, %f124, %f44, %p11;
	selp.f32 	%f126, 0fC3170000, 0fC2FE0000, %p11;
	mov.b32 	%r39, %f125;
	and.b32  	%r40, %r39, 8388607;
	or.b32  	%r41, %r40, 1065353216;
	mov.b32 	%f127, %r41;
	shr.u32 	%r42, %r39, 23;
	cvt.rn.f32.u32 	%f128, %r42;
	add.f32 	%f129, %f126, %f128;
	setp.gt.f32 	%p12, %f127, 0f3FB504F3;
	mul.f32 	%f130, %f127, 0f3F000000;
	add.f32 	%f131, %f129, 0f3F800000;
	selp.f32 	%f132, %f131, %f129, %p12;
	selp.f32 	%f133, %f130, %f127, %p12;
	add.f32 	%f134, %f133, 0fBF800000;
	add.f32 	%f135, %f133, 0f3F800000;
	rcp.approx.ftz.f32 	%f136, %f135;
	add.f32 	%f137, %f134, %f134;
	mul.f32 	%f138, %f137, %f136;
	mul.f32 	%f139, %f138, %f138;
	mov.f32 	%f140, 0f3C4CAF63;
	mov.f32 	%f141, 0f3B18F0FE;
	fma.rn.f32 	%f142, %f141, %f139, %f140;
	mov.f32 	%f143, 0f3DAAAABD;
	fma.rn.f32 	%f144, %f142, %f139, %f143;
	mul.rn.f32 	%f145, %f144, %f139;
	mul.rn.f32 	%f146, %f145, %f138;
	sub.f32 	%f147, %f134, %f138;
	add.f32 	%f148, %f147, %f147;
	neg.f32 	%f149, %f138;
	fma.rn.f32 	%f150, %f149, %f134, %f148;
	mul.rn.f32 	%f151, %f136, %f150;
	add.f32 	%f152, %f146, %f138;
	sub.f32 	%f153, %f138, %f152;
	add.f32 	%f154, %f146, %f153;
	add.f32 	%f155, %f151, %f154;
	add.f32 	%f156, %f152, %f155;
	sub.f32 	%f157, %f152, %f156;
	add.f32 	%f158, %f155, %f157;
	mov.f32 	%f159, 0f3F317200;
	mul.rn.f32 	%f160, %f132, %f159;
	mov.f32 	%f161, 0f35BFBE8E;
	mul.rn.f32 	%f162, %f132, %f161;
	add.f32 	%f163, %f160, %f156;
	sub.f32 	%f164, %f160, %f163;
	add.f32 	%f165, %f156, %f164;
	add.f32 	%f166, %f158, %f165;
	add.f32 	%f167, %f162, %f166;
	add.f32 	%f168, %f163, %f167;
	sub.f32 	%f169, %f163, %f168;
	add.f32 	%f170, %f167, %f169;
	mul.rn.f32 	%f171, %f122, %f168;
	neg.f32 	%f172, %f171;
	fma.rn.f32 	%f173, %f122, %f168, %f172;
	fma.rn.f32 	%f174, %f122, %f170, %f173;
	mov.f32 	%f175, 0f00000000;
	fma.rn.f32 	%f176, %f175, %f168, %f174;
	add.rn.f32 	%f177, %f171, %f176;
	neg.f32 	%f178, %f177;
	add.rn.f32 	%f179, %f171, %f178;
	add.rn.f32 	%f180, %f179, %f176;
	mov.b32 	%r43, %f177;
	setp.eq.s32 	%p13, %r43, 1118925336;
	add.s32 	%r44, %r43, -1;
	mov.b32 	%f181, %r44;
	add.f32 	%f182, %f180, 0f37000000;
	selp.f32 	%f45, %f182, %f180, %p13;
	selp.f32 	%f183, %f181, %f177, %p13;
	mov.f32 	%f184, 0f3FB8AA3B;
	mul.rn.f32 	%f185, %f183, %f184;
	cvt.rzi.f32.f32 	%f186, %f185;
	abs.f32 	%f187, %f186;
	setp.gt.f32 	%p14, %f187, 0f42FC0000;
	mov.b32 	%r45, %f186;
	and.b32  	%r46, %r45, -2147483648;
	or.b32  	%r47, %r46, 1123811328;
	mov.b32 	%f188, %r47;
	selp.f32 	%f189, %f188, %f186, %p14;
	mov.f32 	%f190, 0fBF317218;
	fma.rn.f32 	%f191, %f189, %f190, %f183;
	mov.f32 	%f192, 0f3102E308;
	fma.rn.f32 	%f193, %f189, %f192, %f191;
	mul.f32 	%f194, %f193, 0f3FB8AA3B;
	add.f32 	%f195, %f189, 0f4B40007F;
	mov.b32 	%r48, %f195;
	shl.b32 	%r49, %r48, 23;
	mov.b32 	%f196, %r49;
	ex2.approx.ftz.f32 	%f197, %f194;
	mul.f32 	%f46, %f197, %f196;
	setp.eq.f32 	%p15, %f46, 0f7F800000;
	mov.f32 	%f241, 0f7F800000;
	@%p15 bra 	$L__BB0_10;

	fma.rn.f32 	%f241, %f46, %f45, %f46;

$L__BB0_10:
	mov.f32 	%f233, 0f40000000;
	cvt.rzi.f32.f32 	%f232, %f233;
	add.f32 	%f231, %f232, %f232;
	mov.f32 	%f230, 0f40800000;
	sub.f32 	%f229, %f230, %f231;
	abs.f32 	%f228, %f229;
	setp.lt.f32 	%p16, %f42, 0f00000000;
	setp.eq.f32 	%p17, %f228, 0f3F800000;
	and.pred  	%p1, %p16, %p17;
	setp.eq.f32 	%p18, %f42, 0f00000000;
	@%p18 bra 	$L__BB0_14;
	bra.uni 	$L__BB0_11;

$L__BB0_14:
	add.f32 	%f202, %f42, %f42;
	selp.f32 	%f243, %f202, 0f00000000, %p17;
	bra.uni 	$L__BB0_15;

$L__BB0_11:
	mov.b32 	%r50, %f241;
	xor.b32  	%r51, %r50, -2147483648;
	mov.b32 	%f198, %r51;
	selp.f32 	%f243, %f198, %f241, %p1;
	setp.geu.f32 	%p19, %f42, 0f00000000;
	@%p19 bra 	$L__BB0_15;

	mov.f32 	%f199, 0f40800000;
	cvt.rzi.f32.f32 	%f200, %f199;
	setp.eq.f32 	%p20, %f200, 0f40800000;
	@%p20 bra 	$L__BB0_15;

	mov.f32 	%f243, 0f7FFFFFFF;

$L__BB0_15:
	add.f32 	%f203, %f44, 0f40800000;
	mov.b32 	%r52, %f203;
	setp.lt.s32 	%p22, %r52, 2139095040;
	@%p22 bra 	$L__BB0_20;

	setp.gtu.f32 	%p23, %f44, 0f7F800000;
	@%p23 bra 	$L__BB0_19;
	bra.uni 	$L__BB0_17;

$L__BB0_19:
	add.f32 	%f243, %f42, 0f40800000;
	bra.uni 	$L__BB0_20;

$L__BB0_17:
	setp.neu.f32 	%p24, %f44, 0f7F800000;
	@%p24 bra 	$L__BB0_20;

	selp.f32 	%f243, 0fFF800000, 0f7F800000, %p1;

$L__BB0_20:
	ld.const.u32 	%r65, [params+340];
	mov.f32 	%f204, 0f3F800000;
	sub.f32 	%f205, %f204, %f243;
	setp.eq.f32 	%p25, %f42, 0f3F800000;
	selp.f32 	%f206, 0f00000000, %f205, %p25;
	cvt.sat.f32.f32 	%f244, %f206;
	setp.lt.s32 	%p26, %r65, 2;
	@%p26 bra 	$L__BB0_22;

	setp.lt.f32 	%p27, %f100, 0f00000000;
	neg.f32 	%f207, %f100;
	selp.f32 	%f208, %f207, %f100, %p27;
	mul.f32 	%f244, %f208, %f244;

$L__BB0_22:
	selp.f32 	%f209, 0f00000000, 0f3F800000, %p2;
	mul.f32 	%f210, %f209, %f237;
	mul.f32 	%f211, %f210, %f244;
	mul.f32 	%f212, %f209, %f238;
	mul.f32 	%f213, %f212, %f244;
	mul.f32 	%f214, %f209, %f239;
	mul.f32 	%f215, %f214, %f244;
	mul.f32 	%f216, %f211, 0f437F0000;
	mul.f32 	%f217, %f213, 0f437F0000;
	mul.f32 	%f218, %f215, 0f437F0000;
	cvt.rmi.f32.f32 	%f219, %f216;
	cvt.rmi.f32.f32 	%f220, %f217;
	cvt.rmi.f32.f32 	%f221, %f218;
	mul.f32 	%f222, %f219, %f9;
	mul.f32 	%f223, %f220, %f10;
	mul.f32 	%f224, %f221, %f11;
	setp.gt.f32 	%p29, %f60, %f8;
	selp.f32 	%f225, %f222, %f219, %p29;
	selp.f32 	%f226, %f223, %f220, %p29;
	selp.f32 	%f227, %f224, %f221, %p29;
	mov.b32 	%r54, %f225;
	mov.u32 	%r53, 4;
	// begin inline asm
	call _optix_set_payload, (%r53, %r54);
	// end inline asm
	mov.b32 	%r56, %f226;
	mov.u32 	%r55, 5;
	// begin inline asm
	call _optix_set_payload, (%r55, %r56);
	// end inline asm
	mov.b32 	%r58, %f227;
	mov.u32 	%r57, 6;
	// begin inline asm
	call _optix_set_payload, (%r57, %r58);
	// end inline asm

$L__BB0_24:
	ret;

}

